#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable

void AtomicAdd(volatile __global float *source, const float operand) {
	//union() function can be used to access bit values of a variable
    union {
        unsigned int intVal;
        float floatVal;
    } newVal;
    union {
        unsigned int intVal;
        float floatVal;
    } prevVal;
    do {
        prevVal.floatVal = *source;
        newVal.floatVal = prevVal.floatVal + operand;
    } while (atomic_cmpxchg((volatile __global unsigned int *)source, prevVal.intVal, newVal.intVal) != prevVal.intVal);
}
__kernel
void atomic_test(__global float *answer) {
	int row = get_global_id(0);
	float accumulator = 0;
	for(int i = 0; i < 100000; i++) {
		accumulator = i;
		AtomicAdd(&answer[row], accumulator);
	}
}
